Bloques, índices y un primer ejercicio

Hablaremos de la indexación que normalmente se utiliza en los códigos de CUDA C y que además nos ayudará a comprender mejor los bloques.

Sin más, presentamos el índice idx que, por cierto, se utilizará en el ejercicio a continuación.

idx = blockIdx.x * blockDim.x + threadIdx.x ;

Los componentes de idx son lo suficientemente claros. blockIdx.x se refiere al índice del bloque dentro de una malla, mientras que threadIdx.x se refiere al índice de un thread dentro del block en el que se encuentra. blockDim.x, por otra parte, se refiere a la dimensión del bloque en la dirección x. Sencillo, ¿no?

Estamos suponiendo que únicamente tenemos una dimensión, pero todo esto es análogo en dos o tres dimensiones.

Para ilustrar la magia de este índice, supongamos que tenemos un vector unidimensional con 16 entradas, dividido en 4 bloques de 4 entradas.

Entonces blockDim.x = 4 y es un valor fijo.

blockIdx.x y threadIdx.x van de 0 a 3.

En el primer bloque (blockIdx.x = 0), idx = 0 * 4 + threadIdx.x irá de 0 a 3.

En el segundo (blockIdx.x = 1), idx = 4 + threadIdx.x empezará en 4 y terminará en 7.

Así, sucesivamente hasta terminar en 15 y contando así las 16 entradas de nuestro vector.

Ahora bien. ¿Dónde podemos fijar las dimensiones de los bloques y mallas?

CUDA C ha creado unas variables llamadas dim3 con las que podemos fijar muy sencillamente las dimensiones de estos objetos. Su sintáxis es muy sencilla:

dim3 dimBlock(4, 1, 1) ;
dim3 dimGrid(4, 1, 1) ;

Las variables dimBlock y dimGrid fueron escritas como para el ejemplo pasado. La sintáxis es, como se puede intuir, la ya establecida (x, y, z), por lo que las dos variables recién escritas se refieren a una malla unidimensional (en la dirección x) con 4 bloques. Cada uno de estos últimos también será unidimensional en la dirección x y contará con 4 threads.

Ejercicio: Multiplicación de Vectores

Ahora veremos como escribir nuestro primer código en CUDA C para realizar un ejercicio básico: la multiplicación de vectores.

La dinámica será esta: A continuación pondremos una parte del código. El lector sólo tendrá que llenar las partes faltantes. Prácticamente todos los elementos para poder completar este primer códido están presentes en el notebook anterior, por lo que no hay que dudar en tomarlo como referencia.

Así de fácil.

Nota: las partes que faltan por llenar en este código son

  • el kernel
  • la alocación, copia y liberación de memoria
  • dimensiones de bloques y grid

In [ ]:
%%writefile Programas/Mul_vectores.cu 

__global__ void multiplicar_vectores(float * device_A, float * device_B, float * device_C, int TAMANIO)
{
 // Llena el kernel escribiendo la multiplicacion de los vectores A y B
}

int main( int argc, char * argv[])
{

    int TAMANIO 1000 ;
    float h_A[TAMANIO] ;
    float h_B[TAMANIO] ;
    float h_C[TAMANIO] ;
    float prueba ;
    
    for (int i = 0, i < TAMANIO, i ++)
    {
        h_A[i] = i ;
        h_B[i] = i + 1 ;
    }
    
    // Escribe abajo las lineas para la alocacion de memoria
    
    // Escribe abajo las lineas para copia de memoria del CPU al GPU
    
    // Completa para escribir las dimensiones de los bloques
    
    dim3 dimBlock( ) ;
    dim3 dimGrid( ) ;
    
    // Completa para lanzar el kernel
    
    multiplicar_vectores<<< dimGrid, dimBlock >>>( ) ;
    
    // Copia la memoria del GPU al CPU
    
    
    // Aqui abajo YA ESTAN ESCRITAS las lineas para liberar la memoria
    
    cudaFree(d_A) ;
    cudaFree(d_B) ;
    cudaFree(d_C) ;
    
    
    // Aqui abajo un pequenio codigo para poder saber si tu resultado es correcto
    
    prueba = 0. ;
    
    for (int i = 0, i < TAMANIO, i ++)
    {
        prueba += h_C[i] ;
    }
    
    println( prueba) ;
    return 0;
}

In [ ]:
!nvcc -o Programas/Mul_vectores Programas/Mul_vectores.cu

El resultado final tendría que ser igual a 333300.0. Si tu resultado es correcto, ¡felicidades! has escrito correctamente tu primer código en CUDA C.

En caso de que no logres encontrar el resultado correcto y ya estés harto de intentar, recuerda que nos puedes contactar a nuestros correos (que están en el primer notebook).


In [4]:
suma = 0.
for i in xrange(100):
    suma += i*(i+1)
suma

# Estas cuatro lineas hacen lo mismo que todo el código que hicimos anteriormente
# No se desanimen :(, los verdaderos resultados vendrán muy prontamente.


Out[4]:
333300.0

Este ejercicio tiene distintos objetivos.

  • Hacer notar todos los elementos que un código de CUDA C ha de tener.
  • Encontrar dentro del gran número de elementos que conforma un código de CUDA C, que la escencia de C no se ha ido (por ejemplo, la simplicidad del kernel)
  • Ahondar en el tema de las dimensiones de blocks y grids, como a continuación será descrito.

Un ejercicio muy interesante dentro del que se acaba de realizar es perderle miedo a jugar con los bloques. Esto compilando el código para distintas dimensiones de bloques y mallas.

En principio dos distintas dimensiones pueden dar el mismo resultado correcto, sin embargo más adelante veremos que la elección de una buena dimensión afecta en la optimización del código.

Invitamos al lector a jugar no sólo con las dimensiones ya descritas, sino con otros elementos del código, con el fin que pierda el miedo de escribir lo que se le venga en mente.